#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>  // threadIdx

using namespace std;

// 方式2，声明共享变量，不能给初始值，需要由线程来初始化
__shared__ int sharedValue1;

static __global__ void testKernel() {
    // 方式1，声明静态共享内存变量（静态的意思时编程时确定共享内存大小），block内的线程共用
    __shared__ int sharedArray[8];

    // 方式2，声明共享的变量，不能给初始值，需要由线程来初始化
    __shared__ int sharedValue2;

    if (blockIdx.x == 0 && threadIdx.x == 0) {
        printf("shared var sharedValue1 is inited by thread[%d,%d]\n", blockIdx.x, threadIdx.x);
        sharedValue1 = 10 + blockIdx.x;
        sharedValue2 = 20 + blockIdx.x;
        sharedArray[0] = 30 + blockIdx.x;
    }

    __syncthreads();  // 等待线程块中其他线程都执行到此行，然后再继续往下执行

    printf("thread [%d,%d] -> sharedValue1=%d, sharedValue2=%d, sharedArray[0]=%d, &sharedArray=%ld, &sharedValue1=%ld, &sharedValue2=%ld\n",
            blockIdx.x, threadIdx.x, sharedValue1, sharedValue2, sharedArray[0], &sharedArray, &sharedValue1, &sharedValue2);

}

int main() {
    testKernel<<<2,2>>>();
    cudaDeviceSynchronize();
    return 0;
}

// shared var sharedValue1 is inited by thread[0,0]
// thread [1,0] -> sharedValue1=0, sharedValue2=0, sharedArray[0]=0, &sharedArray=140093394255876, &sharedValue1=140093394255872, &sharedValue2=140093394255908
// thread [1,1] -> sharedValue1=0, sharedValue2=0, sharedArray[0]=0, &sharedArray=140093394255876, &sharedValue1=140093394255872, &sharedValue2=140093394255908
// thread [0,0] -> sharedValue1=10, sharedValue2=20, sharedArray[0]=30, &sharedArray=140093394255876, &sharedValue1=140093394255872, &sharedValue2=140093394255908
// thread [0,1] -> sharedValue1=10, sharedValue2=20, sharedArray[0]=30, &sharedArray=140093394255876, &sharedValue1=140093394255872, &sharedValue2=140093394255908

// 共享内存变量是在一个block内共享的


// __syncthreads函数在CUDA编程中主要用于线程块内的同步‌。它的主要作用是确保在一个线程块中的所有线程在执行到__syncthreads函数之前都已完成它们之前的所有指令。一旦所有线程都到达这个同步点，它们才可以继续执行__syncthreads之后的指令‌
// 使用场景和注意事项
// ‌共享内存的读写‌：当线程写入共享内存，并且这些数据将被线程块中的其他线程所读取时，通常需要一个__syncthreads调用以确保写入完成。这样可以避免竞态条件和数据不一致的问题‌
// ‌避免竞态条件‌：当线程块内的线程可能同时写入同一个位置或在其他线程完成某些操作之前需要读取数据时，使用__syncthreads可以避免这些问题‌
// ‌不要在分支条件下不均匀地调用‌：如果线程块中的一些线程调用了__syncthreads，但其他线程由于某些条件（如if语句）没有调用，可能会导致死锁‌
// ‌不要在循环中过度使用‌：过度使用__syncthreads可能会降低性能，因为它会阻止线程并行地执行‌
// 历史背景和基本概念
// CUDA编程中的同步机制非常重要，因为GPU中的多个线程需要协调执行以避免数据竞争和不一致。__syncthreads函数确保在一个线程块内的所有线程在继续执行之前都已完成它们的操作，从而保证数据的一致性和正确的程序执行顺序‌



// 在CUDA的kernel函数中，使用`extern __shared__`声明的`cache`变量是**每个线程块(block)共享一个**，而不是每个线程(thread)或每个线程束(warp)拥有独立的副本。

// ### 关键点说明：

// 1. **共享内存的作用域**：
//    - 每个线程块拥有自己独立的`cache`副本
//    - 同一个块内的所有线程共享访问同一个`cache`数组
//    - 不同线程块的`cache`是完全独立的

// 2. **内存大小确定**：
//    - 大小是在**启动kernel时动态指定**的（通过第三个执行配置参数）
//    ```c
//    kernel<<<grid_size, block_size, shared_mem_size>>>(...);
//    ```
//    - 例如，指定每个block使用4KB共享内存：
//    ```c
//    kernel<<<128, 256, 4096>>>(...);
//    ```

// 3. **实际数量**：
//    - 等于**启动的线程块数量**（grid维度）
//    - 例如`<<<100, 256, 1024>>>`表示有100个`cache`数组实例（每个block一个）

// 4. **访问特性**：
//    ```c
//    __global__ void kernel(...) {
//        extern __shared__ int cache[]; // 所有线程看到的是自己所在block的cache
//        cache[threadIdx.x] = ...;     // 线程访问自己block内的共享内存
//    }
//    ```

// 5. **容量限制**：
//    - 每个SM的共享内存总量有限（例如Turing架构每个SM最多64KB）
//    - 单个block可用的共享内存大小受限于：
//      - 设备计算能力（常见最大48KB）
//      - 内核启动时实际分配的大小

// ### 示例说明：
// 如果这样启动内核：
// ```c
// kernel<<<100, 256, 4096>>>(...);
// ```
// 则表示：
// - 共100个线程块
// - 每个块有256个线程
// - 每个块有4096字节的共享内存
// - 最终会有100个独立的`cache`数组存在（每个block一个）

// 这种设计使得同一个block内的线程可以通过共享内存高效协作，而不同block之间则完全隔离，确保了线程块级别的并行独立性。